How does OpenCL work?¶
Kernels in software threads¶
An OpenCL implementation is a framework to support running lightweight pieces of code, called kernels, in parallel over the available cores of a compute device. Below is an example kernel to compute the floating point absolute value of a single element in an array of floating point numbers.
__kernel void vec_fabs(
// Memory allocations that are on the compute device
__global float *src,
__global float *dst,
// Number of elements in the memory allocations
int length) {
// Get our position in the array
size_t gid0 = get_global_id(0);
// Get the absolute value of
if (gid0 < length) {
dst[gid0] = fabs(src[gid0]);
}
}
In order to take the absolute value of every element we need to run this kernel at every point in the array. A software thread can be thought of as the execution of a sequence of compute instructions independently from other threads. In that sense a kernel is run in a software thread.
Hardware threads¶
A hardware thread is a pipeline of physical machinery that executes the instructions in a software thread. Compute devices have a number of cores to manage memory and execute software threads. In AMD and OpenCL terminology these cores are called Compute Units. Every compute unit makes available to the OS a number of hardware threads for running kernels. In OpenCL terminology a processing element is synonymous with a hardware thread.
GPU specifics¶
GPU's use a SIMT (Single Instruction Multiple Threads) processing model, where instructions are executed by the Compute Unit over teams of hardware threads that operate in lock-step with each other and in parallel. For AMD GPUs the team is 32-64 threads wide and is called a wavefront. For NVIDIA GPUs a team is usually 32 threads wide and is called a warp. Each hardware thread in a GPU has access to specialised floating point and integer units for perfoming math operations. These specialised units are known as shader cores (AMD) or CUDA cores (NVIDIA). There are many thousands of these units in a GPU and this feature, along with greater memory bandwidth is responsible for the significant performance advantage that GPU's have over CPU's.
The example below a graphical layout of an AMD MI250X GPU processor. Each processor contains two GPU dies; each die contains 8 shader engines; and each shader engine contains ~14 compute units, for a total of 110 Compute Units per die. Every compute unit commands a wavefront of 64 hardware threads, therefore on this processor there are two unique compute devices, each with $110\times64 = 7040$ available hardware threads for use in compute applications.
CPU specifics¶
In CPUs each compute unit also makes available to the OS a number of hardware threads - usually 2-4. These threads are more "independent" than their GPU equivalents and are not constrained to operate in lock-step with each other. CPU hardware threads have access to SIMD vector units to perform vector math operations, however this hardware is only accessed through special vector instructions that the compiler conservatively generates if it deems it is safe to do so. OpenCL has vector data types whose math operations map easily to SIMD vector instructions, thus making it easier to build software that uses SIMD vector units in CPUs.
